/*
 *  Copyright 2010 INFN - APE group
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

#pragma once

#include <cuqu/config.h>

namespace cuqu
{
    namespace detail
    {
        struct barrier_shm {
            atomic_t m_cnt;
        };

#define CUQU_INIT_BARRIER_SHM { 0 }

        static __inline__ __device__ int barrier_wait(barrier_shm *gbl_barrier, int n_waiters) {

            int retcode = SUCCESS;
            const uint tid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
            const uint bid = blockIdx.x + gridDim.x*blockIdx.y;
            //const uint block_size = blockDim.x*blockDim.y*blockDim.z;
            //const uint gridSize = gridDim.y*gridDim.x;
            //const uint gid = tid + bid*block_size;

            // every block leader thread does mod_inc
            if(0 == tid) {
                unsigned int oldcnt = atomic_mod_inc(&gbl_barrier->m_cnt, n_waiters-1);
#if defined(CUQU_DEBUG)
                if(oldcnt >= n_waiters) {
                    cuqu_error("ERROR invalid barrier counter %d\n", oldcnt);
                    return;
                }
#endif
            }
            __syncthreads();

            // every thread in every block loops till the counter reaches 0 due to mod_inc
            int retry = 100000;
            do {
                int cnt;
                __shared__ int shr_cnt;

                //__threadfence();
                if(0 == tid) {
                    shr_cnt = atomic_read(&gbl_barrier->m_cnt);
                    // flush write to shm_cnt
                    __threadfence_block();
                }
                __syncthreads();
                cnt = shr_cnt;
                // not == n_waiters, due to mod_inc
                if(0 == cnt) {
                    cuqu_trace("barrier done, counter reached %d\n", cnt);
                    break;
                }
#if defined(CUQU_DEBUG)
                else if(cnt >= n_waiters) {
                    cuqu_error("ERROR invalid barrier counter %d\n", cnt);
                    break;
                }
#endif
                --retry;
                if(retry < 0) {
                    cuqu_error("%d:%d max retry count exceeded, last cnt=%d\n", bid, tid, cnt);
                    retcode = WOULDBLOCK;
                    break;
                }
            } while(true);
            __syncthreads();
            return retcode;
        }

    } // end namespace detail

} // end namespace cuqu

/*
 * Local variables:
 *  mode: c++
 *  c-indent-level: 4
 *  c-basic-offset: 4
 *  tab-width: 4
 *  indent-tabs-mode: nil
 * End:
 */
